//memoryManage.cu：内存管理
/*Authored by alpc40*/


#include <stdio.h>
#include<stdlib.h>
//下面两个头文件基本都会用到，别省了
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

//这个宏还是蛮有用,主要功能是检测cuda函数的错误，以方便检查
/*
  1.cudaError_t:CUDA的错误码，它的错误码太多了，大概有80-90个，具体什么错误碰到了再解释
  2.cudaGetErrorString():将CUDA错误码转换成具体错误语句，方便判断程序员错误
  3.__FILE__,__LINE__:这个就不解释了，虽然我也用的不太熟
*/
#define CHECK(call)                                                  \
{                                                                    \
    const cudaError_t err = call;                                    \
    if (err != cudaSuccess)                                          \
    {                                                                \
        fprintf(stderr, "Error: %s:%d, ", __FILE__, __LINE__);       \
        fprintf(stderr, "code: %d, reason: %s\n", err,               \
                cudaGetErrorString(err));                            \
        exit(1);                                                     \
    }                                                                \
}



#define NUM 9//数组大小
//内核函数定义
__global__ void sumArrayOnGpu(int* d_arrayA, int* d_arrayB, int* d_result)
{
	int i = blockDim.x * blockIdx.x + threadIdx.x;
	if (i < NUM)//为什么要加这句呢？因为i的取值范围是[0,11](grid*block=12),要养成良好习惯，防止内存溢出
		d_result[i] = d_arrayA[i] + d_arrayB[i];
}
int main4()
{
	int arrayA[NUM] = { 3,4,6,7,8,9,1,19,35 };
	int arrayB[NUM] = { 1,4,7,4,2,8,2,20,33 };
	int result[NUM];//记录CPU结果
	int* d_arrayA, * d_arrayB, * d_result;
	int nBytes = NUM * sizeof(int);//注意拷贝的是字节数，不是个数
	CHECK(cudaMalloc((int**)&d_arrayA, nBytes));//如果觉得CHECK比较烦人，而你不认为这个函数会出错，那么删掉又何妨
	CHECK(cudaMalloc((int**)&d_arrayB, nBytes));//注意申请GPU内存的方法
	CHECK(cudaMalloc((int**)&d_result, nBytes));
	CHECK(cudaMemcpy(d_arrayA, arrayA, nBytes, cudaMemcpyHostToDevice));//注意拷贝的方法
	CHECK(cudaMemcpy(d_arrayB, arrayB, nBytes, cudaMemcpyHostToDevice));
	int thd_num = 4;
	dim3 block(thd_num);//block与grid的定义，结果为(4,1,1)，即每个内存块有线程4个
	dim3 grid((NUM + thd_num - 1) / thd_num);//结果为(3,1,1)，即内存块为3个
	sumArrayOnGpu << <grid, block >> > (d_arrayA, d_arrayB, d_result);//内核函数
	CHECK(cudaMemcpy(result, d_result, nBytes, cudaMemcpyDeviceToHost));
	printf("计算结果：\n");
	for (int i = 0; i < NUM; i++)
		printf(" %d", result[i]);
	printf("\n");
	cudaFree(d_arrayA);//回收GPU内存，一般申请时立马写回收函数，防止忘记导致内存泄露
	cudaFree(d_arrayB);
	cudaFree(d_result);
	cudaDeviceReset();//GPU恢复
	return 0;
}